/**
	multi-column sorting
*/

#ifndef __MC_SORT_CU__
#define __MC_SORT_CU__

#include <stdio.h>
#include <stdlib.h>
#include "cuda_header.h"
#include <thrust/sequence.h>
#include <thrust/sort.h>

//struct __align__(8) st_data {
struct st_data {
	unsigned int data1;
	unsigned char data2;
	unsigned char data3;
	unsigned char data4;
	unsigned char data5;
	unsigned char data6;

	void randNum() {
		data1 = rand();
		data2 = rand()%256;
		data3 = rand()%256;
		data4 = rand()%256;
		data5 = rand()%256;
		data6 = rand()%256;
	}

	void print() {
		printf("%u, %d, %d, %d, %d, %d", data1, data2, data3, data4, data5, data6);
	}
};

struct st_datacpu {
	unsigned int key;
	st_data data;
};

int compare(const void* a, const void* b) {
	return ((st_datacpu*)a)->key - ((st_datacpu*)b)->key;
}

void mcsort(unsigned int* d_key, st_data* d_dataout, st_data* d_datain,
		const unsigned int numElement) {

	printf("sizeof(st_data) = %d\n", sizeof(st_data));

	//make the identity
	unsigned int* d_id = NULL;
	GPUMALLOC((void**)&d_id, sizeof(unsigned int)*numElement);
	thrust::device_ptr<unsigned int> dptr_id(d_id);	
	thrust::sequence(dptr_id, dptr_id + numElement);	
	unsigned int timer = 0;

	//sorting
	//timer = 0;
	//startTimer(&timer);
	thrust::device_ptr<unsigned int> dptr_key(d_key);
	thrust::sort_by_key(dptr_key, dptr_key + numElement, dptr_id);
	cutilSafeCall(cudaThreadSynchronize());
	//endTimer(&timer, "sort_by_key");

	//gather
	//timer = 0;
	//startTimer(&timer);
	thrust::device_ptr<st_data> dptr_datain(d_datain);
	thrust::device_ptr<st_data> dptr_dataout(d_dataout);
	thrust::gather(dptr_id, dptr_id + numElement, dptr_datain, dptr_dataout);
	cutilSafeCall(cudaThreadSynchronize());
	//endTimer(&timer, "gather");

	//memory cleanup
	GPUFREE(d_id);
}

/*
//for debug
int main(int argc, char** argv) {

	const unsigned int numElement = 123*1024*1024;
	printf("numElement = %u\n", numElement);
	unsigned int* h_key = (unsigned int*)malloc(sizeof(unsigned int)*numElement);
	st_data* h_data = (st_data*)malloc(sizeof(st_data)*numElement);
	for(unsigned int i = 0; i < numElement; i++) {
		h_key[i] = rand();
		h_data[i].randNum();
	}
	unsigned int* d_key = NULL;
	GPUMALLOC((void**)&d_key, sizeof(unsigned int)*numElement);
	TOGPU(d_key, h_key, sizeof(unsigned int)*numElement);
	st_data* d_datain = NULL;
	GPUMALLOC((void**)&d_datain, sizeof(st_data)*numElement);
	TOGPU(d_datain, h_data, sizeof(st_data)*numElement);
	st_data* d_dataout = NULL;
	GPUMALLOC((void**)&d_dataout, sizeof(st_data)*numElement);


	unsigned int timer = 0;
	startTimer(&timer);
	mcsort(d_key, d_dataout, d_datain, numElement);
	cutilSafeCall(cudaThreadSynchronize());
	endTimer(&timer, "mcsort");

	//sorting on the CPU
	st_datacpu* gold = (st_datacpu*)malloc(sizeof(st_datacpu)*numElement);
	for(unsigned int i = 0; i < numElement; i++) {
		gold[i].key = h_key[i];
		gold[i].data = h_data[i];
	}
	timer = 0;
	startTimer(&timer);
	qsort(gold, numElement, sizeof(st_datacpu), compare);
	endTimer(&timer, "qsort");	


	//copy result back
	FROMGPU(h_key, d_key, sizeof(unsigned int)*numElement);
	FROMGPU(h_data, d_dataout, sizeof(st_data)*numElement);


	//check result
	for(unsigned int i = 0; i < numElement; i++) {
		if(h_key[i] != gold[i].key ||
		h_data[i].data1 != gold[i].data.data1||
		h_data[i].data2 != gold[i].data.data2||
		h_data[i].data3 != gold[i].data.data3||
		h_data[i].data4 != gold[i].data.data4||
		h_data[i].data5 != gold[i].data.data5||
		h_data[i].data6 != gold[i].data.data6) {
			printf("error!!!\n");
		}
	}
	printf("Test pass\n");

	GPUFREE(d_key);
	GPUFREE(d_datain);
	GPUFREE(d_dataout);

	return EXIT_SUCCESS;
}
*/
#endif /*__MC_SORT_CU__*/


